#include "CudaTexture.cuh"
#include <algorithm/rayTracing/qmcrt_util/check_cuda.h>
#include <algorithm/shared/tool/imageSaver/ImageSaver.h>

#include <iostream>

// #include <sys/time.h>

using namespace solar::cuda;

// static constexpr int dim = 1000;

// static auto cpuSecond() -> double
// {
//     struct timeval tp;
//     gettimeofday(&tp, NULL);
//     return ((double)tp.tv_sec + (double)tp.tv_usec * 1e-6);
// }

/*
__device__ float clamp(float x, float a, float b)
{
    return max(a, min(b,x));
}

__device__ float4 rgbToInt(float r, float g, float b)
{
    r = clamp(r, 0.0f, 255.0f);
    g = clamp(g, 0.0f, 255.0f);
    b = clamp(b, 0.0f, 255.0f);
    printf("r=%f, g=%f, b=%f\n", r, g, b);
    return make_float4(r,g,b,125);
}

static __global__ void cudaProcess(uchar4 *g_data, int imgw)
{
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int bw = blockDim.x;
    int bh = blockDim.y;
    int x = blockIdx.x * bw + tx;
    int y = blockIdx.y * bh + ty;

    float fx = x / (float)DIM - 0.5f;
    float fy = y / (float)DIM - 0.5f ;
    int green = 128+127 * tan(abs(fx*100) - abs(fy*100));
    // uchar4 c4 = make_uchar4((x & 0x20) ? 100 : 0, 0, (y & 0x20) ? 100 : 0, 0);
    // float4 c4 = make_float4(255,110,200,0);
    int idx = y*imgw + x;
    g_data[idx] = make_uchar4(255, 0, 0, 255);
    // g_data[idx] = rgbToInt(c4.z, c4.y, c4.x);
    // g_data[idx] = make_float4((uchar)255.0,(uchar)110.0,(uchar)200.0,(uchar)255.0);
    printf("data[%d]: %f %f %f\n",idx, g_data[idx].x, g_data[idx].y, g_data[idx].z);
}

static __global__ void cudaGetMaxValue(float* data, float* max_value, int n)
{
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx >= n) return;

    float *idata = data + blockIdx.x * blockDim.x;
    for(int stride = blockDim.x / 2 ; stride > 0 ; stride /=2)
    {
        if(tid < stride)
        {
            idata[tid] = max(idata[tid], idata[tid + stride]);
        }
        __syncthreads();
    }
    if(tid==0) max_value[blockIdx.x] = idata[0];

}

__device__ uchar4 getColorAtValueCuda(int3* colors, float* factors, float factor, int color_num)
{
    int3 color;
    int id=0;
    while(factor > factors[id] && id < color_num)
    {
        id++;
    }
    // printf("id=%d, factor=%f\n",id, factor);

    if(id == color_num) color = colors[color_num-1];
    else if(id == 0) color = colors[0];
    else{
        int3 left = colors[id-1];
        int3 right = colors[id];
        float coef = (factor - factors[id-1]) / (factors[id] - factors[id-1]);
        float r = left.x + coef * (right.x - left.x);
        float g = left.y + coef * (right.y - left.y);
        float b = left.z + coef * (right.z - left.z);
        color = make_int3((int)r, (int)g, (int)b);
        // if((color.x==0 && color.y==0&& color.z==0) || color.x > 255 || color.y > 255 || color.z >
255)
            // printf("factor=%f,id=%d,left=(%f, %d,%d,%d), right=(%f, %d,%d,%d),
color=(%d,%d,%d)\n",
            //      factor, id,        factors[id-1], left.x,left.y,left.z,
            //                         factors[id], right.x, right.y, right.z,
            //                                     color.x, color.y, color.z);
    }

    return make_uchar4(color.x , color.y , color.z , 255);
}
*/
/**
 * GPU中：图像rgba、flux、分辨率、colorbar（colorbar已经设置过最大最小值）
 */
/*
static __global__ void cudaGetFluxMap(uchar4 *img, float* data, int width, int height,
                                        int3* colors, float* factors, int color_num, float
max_value)
{
    // printf("1\n");
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int bw = blockDim.x;
    int bh = blockDim.y;
    int x = blockIdx.x * bw + tx;
    int y = blockIdx.y * bh + ty;
    // printf("x=%d, y=%d, width=%d, height=%d\n", x,y, width, height);

    if(x >= width || y >= height) return;


    int idx = y * width + x;
    // printf("x=%d, y=%d, idx=%d\n",x,y,idx);
    img[idx] = getColorAtValueCuda(colors, factors, data[idx]/max_value, color_num);
    // int r = colorbar.getRedComponentAtValue(data[idx]);
    // int g = colorbar.getGreenComponentAtValue(data[idx]);
    // int b = colorbar.getBlueComponentAtValue(data[idx]);

    // img[idx] = make_float4((float)r / 255.0, (float)g / 255.0, (float)b / 255.0, 1.0);
    // if(img[idx].x<1e-2 && img[idx].y<1e-2 && img[idx].z<1e-2)
    // printf("data=%f, max_value=%f, width=%d, height=%d,img[%d]: %f %f %f\n",
    //         data[idx], max_value,width, height, idx, img[idx].x, img[idx].y, img[idx].z);
    // int width, int height
}
*/

/*
    texture---需要新生成的部分，fluxmap_data----raytracingpipeline::result_data_中传来的数据
*/
static __global__ void cudaSetFluxMapTexture(uchar4* texture, uchar4* fluxmap_data, int2 resolution)
{
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int bw = blockDim.x;
    int bh = blockDim.y;
    int x = blockIdx.x * bw + tx;
    int y = blockIdx.y * bh + ty;
    // printf("x=%d, y=%d, width=%d, height=%d\n", x,y, width, height);
    int width = resolution.x;
    int height = resolution.y;
    if (x >= width || y >= height)
    {
        return;
    }

    int idx = y * width + x;
    // printf("x=%d, y=%d, idx=%d\n",x,y,idx);
    texture[idx] = fluxmap_data[idx];
}

/*
void launch_cudaProcess(bool usecuda, dim3 grid, dim3 block, uchar4 *g_data,  float* data, int
width, int height, int3* colors, float* factors, int color_num, float max_value)
{
    printf("launch_cudaprocess--begin\n");
    if(! usecuda )  //若data为nullptr，则生成默认的纹理
    {
        printf("cudaProcess-begin\n");
        cudaProcess<< <grid, block>> > (g_data, DIM);
        cudaDeviceSynchronize();
        printf("cudaProcess-end\n");
    }
    else{       // 否则，按照data的值生成
        printf("cudaGetFluxMap-begin\n");

        auto start_time = std::chrono::high_resolution_clock::now();
        auto end_time = start_time;
        long long elapsed;

        cudaGetFluxMap<<<grid, block>>> (g_data, data, width, height, colors, factors, color_num,
max_value); cudaDeviceSynchronize();

        end_time = std::chrono::high_resolution_clock::now();
        elapsed = std::chrono::duration_cast<std::chrono::microseconds>(end_time -
start_time).count(); printf("\t cuda get flux map took %lld microscends\n", elapsed);
        printf("cudaGetFluxMap<<<grid(%d,%d), block(%d,%d)>>>-end\n", grid.x, grid.y, block.x,
block.y);
    }
    printf("launch_cudaprocess--end\n");

}
*/

/**
--------------------------------------------------------------------------------
*/
// CudaTexture::CudaTexture()
// {

// }

CudaTexture::CudaTexture() : texture_(nullptr), fluxmap_data_(nullptr), resolution_({1000, 1000}) {}
CudaTexture::~CudaTexture()
{
    // checkCudaErrors(cudaFree(colors_));
    // checkCudaErrors(cudaFree(factors_));
}
// void CudaTexture::setColor()
// {
//     colorbar_ = new CudaColorBar();
//     colorbar_->setColor();
//     colorbar_->test();

//     color_num_ = colorbar_->getColorNum();
//     checkCudaErrors(cudaMalloc((void**)&colors_, (color_num_)*sizeof(int3)));
//     checkCudaErrors(cudaMalloc((void**)&factors_, (color_num_)*sizeof(float)));

//     checkCudaErrors(cudaMemcpy((void*)colors_, (void*)colorbar_->getColors(),
//                                 color_num_ * sizeof(int3), cudaMemcpyHostToDevice));
//     checkCudaErrors(cudaMemcpy((void*)factors_, (void*)colorbar_->getFactors(),
//                                 color_num_ * sizeof(float), cudaMemcpyHostToDevice));
// }

// void CudaTexture::calculateMaxValue()
// {
//     // 1D网格，1D线程块
//     dim3 block(512,1);
//     dim3 grid((resolution_.x * resolution_.y + block.x - 1) / block.x, 1);

//     float* tmp;
//     float* max_value_block;
//     cudaMalloc((void**)&tmp, resolution_.x * resolution_.y * sizeof(float));
//     cudaMalloc((void**)&max_value_block, grid.x * sizeof(float));
//     cudaMemcpy(tmp, data_, resolution_.x * resolution_.y * sizeof(float),
//     cudaMemcpyDeviceToDevice); float start_time, elapsed_time; start_time = cpuSecond();
//     cudaGetMaxValue<<<grid, block>>>(tmp, max_value_block, resolution_.x * resolution_.y);
//     cudaDeviceSynchronize();

//     elapsed_time = cpuSecond() - start_time;
//     max_value_=0;
//     float* h_data;
//     h_data = (float*)malloc(grid.x * sizeof(float));
//     cudaMemcpy(h_data, max_value_block, grid.x * sizeof(float), cudaMemcpyDeviceToHost);
//     for(int i=0;i<grid.x;i++) max_value_ = max(max_value_, h_data[i]);

// }
// void CudaTexture::generateTexture()
// {
//     initialTexture();

//     // 注册显卡资源
//     checkCudaErrors(cudaGraphicsGLRegisterImage(&resource_,
//                                                 // texture_->textureId(),
//                                                 texid_,
//                                                 GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone));
//     // memset(&resource_, 0, sizeof(resource_));
//     // resource_.resType = cudaResourceTypeArray;

//     // 分配显卡空间
//     checkCudaErrors(cudaMalloc((void**)&cuPtr_,
//                                 DIM*DIM*sizeof(uchar4)));

//     // 调用核函数写纹理
//     runCuda(false);
//     // writeTextureToFile(texid_, DIM*DIM);
// }

void CudaTexture::generateTextureViaRayTracingPipeline()
{
    printf("CudaTexture::generate Texture Via RayTracingPipeline  --begin\n");

    initialTexture();

    // 注册显卡资源
    checkCudaErrors(cudaGraphicsGLRegisterImage(&resource_,
                                                // texture_->textureId(),
                                                texid_, GL_TEXTURE_2D,
                                                cudaGraphicsRegisterFlagsNone));
    // memset(&resource_, 0, sizeof(resource_));
    // resource_.resType = cudaResourceTypeArray;

    // 分配显卡空间
    checkCudaErrors(cudaMalloc((void**)&texture_, resolution_.x * resolution_.y * sizeof(uchar4)));
    runCuda(true);
    glBindTexture(GL_TEXTURE_2D, 0);
    printf("CudaTexture::generate Texture Via RayTracingPipeline  --end\n");
    calculationFinished = true;
}

void CudaTexture::initialTexture()
{
    /*
    texture_ = new QOpenGLTexture(QOpenGLTexture::Target2D);
    texture_->setMagnificationFilter(QOpenGLTexture::Linear);
    texture_->setMinificationFilter(QOpenGLTexture::Nearest);
    texture_->setWrapMode(QOpenGLTexture::Repeat);
    texture_->setFormat(QOpenGLTexture::RGBA32F);    //格式，需要与cuPtr的类型匹配
    texture_->setSize(resolution_.x, resolution_.y);        //纹理大小与cuPtr的内存大小一致
    texture_->allocateStorage();
*/
    // texture_->setBorderColor(QColor(255, 255,255.255));

    glEnable(GL_TEXTURE_2D);
    glGenTextures(1, &texid_);

    glBindTexture(GL_TEXTURE_2D, texid_);
    glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, resolution_.x, resolution_.y, 0, GL_RGBA,
                 GL_UNSIGNED_BYTE, nullptr);

    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER);

    // glBindTexture(GL_TEXTURE_2D, 0);    //解绑

    std::cout << "=================  " << texid_ << std::endl;
}

void CudaTexture::runCuda(bool usecuda)
{
    cudaArray_t glArray;

    // 映射已经注册的显卡资源到cuda流中
    checkCudaErrors(cudaGraphicsMapResources(1, &resource_, nullptr));

    // resource_.res.array.array = glArray;

    /**
    TODO::get blocks and grids from raytracing pipeline
    */
    dim3 blocks(32, 32);
    dim3 grids((resolution_.x + blocks.x - 1) / blocks.x,
               (resolution_.y + blocks.y - 1) / blocks.y);
    // launch_cudaProcess(usecuda, grids, blocks, cuPtr_, data_, resolution_.x, resolution_.y,
    // colors_, factors_, color_num_, max_value_);
    cudaSetFluxMapTexture<<<grids, blocks>>>(texture_, fluxmap_data_, resolution_);

    // 获取一个数组，通过该数组可以访问已映射图形资源的子资源
    checkCudaErrors(cudaGraphicsSubResourceGetMappedArray(&glArray, resource_, 0, 0));
    printf("cudaMemcpyToArray size %d %d %d\n", resolution_.x, resolution_.y, sizeof(uchar4));

    // ！！！这句如果删去会报错，但是这一函数本身是deprecated的
    checkCudaErrors(cudaMemcpyToArray(glArray, 0, 0, texture_,
                                      resolution_.x * resolution_.y * (int)sizeof(uchar4),
                                      cudaMemcpyDeviceToDevice));
    checkCudaErrors(cudaGraphicsUnmapResources(1, &resource_, nullptr));
    // cudaGraphicsUnmapResources(1, &resource_, 0);
}
// void CudaTexture::setCudaPara(float* data, int width, int height, float max_value)
// {
//     resolution_ = make_int2(width, height);
//     max_value_ = max_value;
//     data_ = data;
//     setColor();
// }
void CudaTexture::setCudaPara(uchar4* fluxmap_data, int2 resolution)
{
    fluxmap_data_ = fluxmap_data;
    resolution_ = resolution;
}

// void CudaTexture::generateTextureViaCudaDataPtr()
// {
//     printf("CudaTexture::generate Texture Via Cuda Data Ptr  --begin\n");

//     initialTexture();

//     // 注册显卡资源
//     checkCudaErrors(cudaGraphicsGLRegisterImage(&resource_,
//                                                 // texture_->textureId(),
//                                                 texid_,
//                                                 GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone));
//     // memset(&resource_, 0, sizeof(resource_));
//     // resource_.resType = cudaResourceTypeArray;

//     // 分配显卡空间
//     checkCudaErrors(cudaMalloc((void**)&cuPtr_,
//                                 resolution_.x*resolution_.y*sizeof(uchar4)));
//     runCuda(true );
//     glBindTexture(GL_TEXTURE_2D, 0);
//     printf("CudaTexture::generate Texture Via Cuda Data Ptr  --end\n");
//     calculationFinished = true;
// }

// 输出 cuPtr中的结果
void printCudaPtr() {}

void CudaTexture::cleanup()
{
    // checkCudaErrors(cudaFree(cuPtr_));

    // checkCudaErrors(cudaGraphicsUnregisterResource(resource_));
}

// void writeTextureToFile(GLuint texture, int size)
// {
//     uchar4* data = new uchar4[size];
//     glGetTextureImage(texture, 0, GL_RGBA, GL_FLOAT, sizeof(uchar4)*size);

//     // access the data
//     std::ofstream myfile("test.txt");
//     if (myfile.is_open())
//     {
//         for (int i = 0; i < size; i++)
//         {
//             myfile << data[i].x << ", " << data[i].y << ", " << data[i].z << ", " << data[i].w <<
//             "\n";
//         }
//         myfile.close();
//     }
//     else
//         std::cout << "Unable to open file" << std::endl;
// }
